/*
 *  Copyright 2010 INFN - APE group
 *
 *  Licensed under the Apache License, Version 2.0 (the "License");
 *  you may not use this file except in compliance with the License.
 *  You may obtain a copy of the License at
 *
 *      http://www.apache.org/licenses/LICENSE-2.0
 *
 *  Unless required by applicable law or agreed to in writing, software
 *  distributed under the License is distributed on an "AS IS" BASIS,
 *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 *  See the License for the specific language governing permissions and
 *  limitations under the License.
 */

#pragma once

#include <cuqu/config.h>
#error "this is superseded by barrier_oneloop.h"

namespace cuqu
{
    namespace detail
    {
        struct barrier_simple {
            atomic_t m_cnt;
            uint     m_last_goal;
        };

        //typedef struct barrier barrier_t;

#define CUQU_INIT_BARRIER_SIMPLE { 0, 0 }

        // WARNING: there may still be some bug here!!!

        static __inline__ __device__ int barrier_wait(barrier_simple *gbl_barrier, int n_waiters)
        {
            const uint tid = threadIdx.x + blockDim.x*threadIdx.y + blockDim.x*blockDim.y*threadIdx.z;
            const uint bid = blockIdx.x + gridDim.x*blockIdx.y;
            const uint block_size = blockDim.x*blockDim.y*blockDim.z;
            const uint grid_size = gridDim.y*gridDim.x;
            const uint gid = tid + bid*block_size;
            int retry = 20000;  
            int retcode = SUCCESS;
            const int cnt_max = 100;
            int cnt = cnt_max;
            uint goal;

            // is it necessary ???
            //__threadfence();

            goal = *(volatile uint*)&gbl_barrier->m_last_goal + n_waiters;

            if(0 == tid) {
                atomic_add(&gbl_barrier->m_cnt, 1);
                while(atomic_read(&gbl_barrier->m_cnt) != goal) {
                    --retry;
                    if(retry < 0) {
                        cuqu_error("[%d:%d] timeout polling, cnt=%d\n", bid, tid, cnt);
                        retcode = WOULDBLOCK;
                        break;
                    }
                };
                if(gid == 0) {
                    gbl_barrier->m_last_goal = goal;
                    // make it visible to all threads in the device for global memory accesses
                    __threadfence();
                    // make it visible to all threads in the device for global memory accesses
                    // and host threads for page-locked host memory accesses
                    //__threadfence_system();
                }
            }
            __syncthreads();
            return retcode;
        }

    } // end namespace detail

} // end namespace cuqu

/*
 * Local variables:
 *  mode: c++
 *  c-indent-level: 4
 *  c-basic-offset: 4
 *  tab-width: 4
 *  indent-tabs-mode: nil
 * End:
 */
